Skip to main content

NVGPUOps

nvgpu.cluster_arrive (triton::nvgpu::ClusterArriveOp)

语法:

operation ::= `nvgpu.cluster_arrive` attr-dict

属性:

属性MLIR 类型描述
relaxed::mlir::IntegerAttr1-bit 无符号整数属性

nvgpu.cluster_id (triton::nvgpu::ClusterCTAIdOp)

语法:

operation ::= `nvgpu.cluster_id` attr-dict

特征:AlwaysSpeculatableImplTrait

接口:ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果:MemoryEffects::Effect{}

结果:

结果描述
result32-bit 无符号整数

nvgpu.cluster_wait (triton::nvgpu::ClusterWaitOp)

语法:

operation ::= `nvgpu.cluster_wait` attr-dict

nvgpu.fence_async_shared (triton::nvgpu::FenceAsyncSharedOp)

语法:

operation ::= `nvgpu.fence_async_shared` attr-dict

属性:

属性MLIR 类型描述
bCluster::mlir::BoolAttr布尔属性

nvgpu.stmatrix (triton::nvgpu::StoreMatrixOp)

语法:

operation ::= `nvgpu.stmatrix` operands attr-dict `:` type(operands)

接口:MemoryEffectOpInterface (MemoryEffectOpInterface)

效果:MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

操作:

操作描述
addr地址空间 3 中的 LLVM 指针
datas32-bit 无符号整数的可变参数

nvgpu.wgmma_commit_group (triton::nvgpu::WGMMACommitGroupOp)

语法:

operation ::= `nvgpu.wgmma_commit_group` attr-dict

nvgpu.wgmma_fence (triton::nvgpu::WGMMAFenceOp)

语法:

operation ::= `nvgpu.wgmma_fence` attr-dict

nvgpu.wgmma (triton::nvgpu::WGMMAOp)

语法:

operation ::= `nvgpu.wgmma` $opA `,` $opB (`,` $opC^)? attr-dict `:` functional-type(operands, $res)

Attributes: 属性:

属性MLIR 类型描述
m::mlir::IntegerAttr32-bit 无符号整数属性
n::mlir::IntegerAttr32-bit 无符号整数属性
k::mlir::IntegerAttr32-bit 无符号整数属性
eltTypeC::mlir::triton::nvgpu::WGMMAEltTypeAttrwgmma 操作数类型为 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', 或者 'f32'
eltTypeA::mlir::triton::nvgpu::WGMMAEltTypeAttrwgmma 操作数类型为 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', 或者 'f32'
eltTypeB::mlir::triton::nvgpu::WGMMAEltTypeAttrwgmma 操作数类型为 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', 或者 'f32'
layoutA::mlir::triton::nvgpu::WGMMALayoutAttrwgmma 布局,可以为 'row'(行)或 'col'(列)
layoutB::mlir::triton::nvgpu::WGMMALayoutAttrwgmma 布局,可以为 'row'(行)或 'col'(列)

操作:

操作描述
opAwgmma operand A/B type wgmma 操作 A/B 类型
opBwgmma operand A/B type wgmma 操作 A/B 类型
opCLLVM structure type LLVM 结构体类型

结果:

结果描述
resLLVM 结构体类型

nvgpu.wgmma_wait_group (triton::nvgpu::WGMMAWaitGroupOp)

语法:

operation ::= `nvgpu.wgmma_wait_group` $input attr-dict `:` type($input)

接口:InferTypeOpInterface

属性:

属性MLIR 类型描述
pendings::mlir::IntegerAttr32-bit 无符号整数属性

操作:

操作描述
inputLLVM structure type LLVM 结构体类型

结果:

结果描述
outputLLVM structure type LLVM 结构体类型